/*
	.struct HeapChunk
	{
		.u64 Next;
		.u16 Magic;
	};
	.struct Heap
	{
		.u32 ChunkCount;
		.u64 FreeChunkChain;
	};

	.global .u64 s_heap;

.entry xtorAlloc
{
	.reg .u32 %r<4>;
	.reg .u64 %rd<9>;
$LBB1_xtorAlloc:
	.loc	3	42	0
	ld.global.u64 	%rd1, [s_heap];
	ld.global.u64 	%rd2, [%rd1+8];
	.loc	3	46	0
	ld.global.u64 	%rd3, [%rd2+0];
	st.global.u64 	[%rd1+8], %rd3;
	.loc	3	47	0
	mov.u64 	%rd4, 0;
	st.global.u64 	[%rd2+0], %rd4;
	.loc	4	35	0
	mov.u64 	%rd5, 0;
	st.global.u64 	[%rd2+16], %rd5;
	.loc	4	36	0
	mov.u64 	%rd6, 0;
	st.global.u64 	[%rd2+24], %rd6;
	.loc	4	37	0
	mov.u32 	%r1, 40;
	.loc	4	38	0
	mov.u32 	%r2, 30806;
	st.global.v2.u16 	[%rd2+32], {%r1,%r2};
	.loc	4	39	0
	add.u64 	%rd7, %rd2, 16;
	st.global.u64 	[%rd2+40], %rd7;
	.loc	4	40	0
	st.global.u64 	[%rd2+48], %rd7;
	.loc	18	15	0
	exit;
$LDWend_xtorAlloc:
	} // xtorAlloc

	.entry AllocDispose (
		.param .u64 __cudaparm_AllocDispose_t)
	{
	.reg .u64 %rd<8>;
	.reg .pred %p<4>;
	.loc	18	16	0
$LBB1_AllocDispose:
	.loc	4	50	0
	ld.param.u64 	%rd1, [__cudaparm_AllocDispose_t];
	ld.global.u64 	%rd2, [%rd1+24];
	mov.u64 	%rd3, 0;
	setp.eq.u64 	%p1, %rd2, %rd3;
	@%p1 bra 	$Lt_1_2306;
	ld.global.u64 	%rd4, [s_heap];
	ld.global.u64 	%rd5, [%rd4+8];
$Lt_1_2818:
 //<loop> Loop body line 58
	.loc	3	58	0
	st.global.u64 	[%rd2+-16], %rd5;
	.loc	3	59	0
	sub.u64 	%rd5, %rd2, 16;
	st.global.u64 	[%rd4+8], %rd5;
	.loc	4	50	0
	ld.global.u64 	%rd2, [%rd2+0];
	mov.u64 	%rd6, 0;
	setp.ne.u64 	%p2, %rd2, %rd6;
	@%p2 bra 	$Lt_1_2818;
$Lt_1_2306:
	.loc	18	16	0
	exit;
$LDWend_AllocDispose:
	} // AllocDispose

	.entry AllocGet (
		.param .u64 __cudaparm_AllocGet_t,
		.param .u16 __cudaparm_AllocGet_bytes,
		.param .s8 __cudaparm_AllocGet_zeroAlloc)
	{
	.reg .u16 %rh<7>;
	.reg .u32 %r<18>;
	.reg .u64 %rd<17>;
	.reg .pred %p<8>;
	.loc	18	17	0
$LBB1_AllocGet:
	.loc	4	64	0
	ld.param.u64 	%rd1, [__cudaparm_AllocGet_t];
	mov.s64 	%rd2, %rd1;
	ld.global.u64 	%rd3, [%rd1+32];
	mov.u64 	%rd4, 0;
	setp.eq.u64 	%p1, %rd3, %rd4;
	@%p1 bra 	$Lt_2_8962;
	ld.param.u16 	%rh1, [__cudaparm_AllocGet_bytes];
$Lt_2_6914:
 //<loop> Loop body line 65
	.loc	4	65	0
	ld.global.u16 	%rh2, [%rd3+16];
	.loc	4	64	0
	ld.param.u16 	%rh1, [__cudaparm_AllocGet_bytes];
	.loc	4	65	0
	add.u16 	%rh3, %rh2, %rh1;
	cvt.u32.u16 	%r1, %rh3;
	cvt.u16.u32 	%r2, %r1;
	mov.u32 	%r3, 1024;
	setp.le.u32 	%p2, %r2, %r3;
	selp.s32 	%r4, 1, 0, %p2;
	cvt.u8.u32 	%r5, %r4;
	.loc	4	66	0
	@%p2 bra 	$Lt_2_9218;
 //<loop> Part of loop body line 65, head labeled $Lt_2_6914
	.loc	4	64	0
	mov.s64 	%rd2, %rd3;
	ld.global.u64 	%rd3, [%rd3+8];
	mov.u64 	%rd5, 0;
	setp.ne.u64 	%p3, %rd3, %rd5;
	@%p3 bra 	$Lt_2_6914;
	cvt.u16.u32 	%rh4, %r2;
	bra.uni 	$Lt_2_258;
$Lt_2_8962:
	cvt.u32.u8 	%r5, %rh5;
	bra.uni 	$Lt_2_258;
$Lt_2_9218:
	cvt.u16.u32 	%rh4, %r2;
$Lt_2_258:
	.loc	4	66	0
	mov.u32 	%r6, 0;
	set.eq.u32.u32 	%r7, %r5, %r6;
	neg.s32 	%r8, %r7;
	mov.u64 	%rd6, 0;
	set.eq.u32.u64 	%r9, %rd3, %rd6;
	neg.s32 	%r10, %r9;
	or.b32 	%r11, %r8, %r10;
	mov.u32 	%r12, 0;
	setp.eq.s32 	%p4, %r11, %r12;
	@%p4 bra 	$Lt_2_7682;
	.loc	3	42	0
	ld.global.u64 	%rd7, [s_heap];
	ld.global.u64 	%rd8, [%rd7+8];
	.loc	3	46	0
	ld.global.u64 	%rd9, [%rd8+0];
	st.global.u64 	[%rd7+8], %rd9;
	.loc	3	47	0
	mov.u64 	%rd10, 0;
	st.global.u64 	[%rd8+0], %rd10;
	.loc	4	70	0
	add.u64 	%rd11, %rd8, 16;
	mov.s64 	%rd3, %rd11;
	.loc	4	64	0
	ld.param.u64 	%rd1, [__cudaparm_AllocGet_t];
	.loc	4	71	0
	ld.global.u64 	%rd12, [%rd1+24];
	st.global.u64 	[%rd8+16], %rd12;
	.loc	4	72	0
	ld.global.u64 	%rd13, [%rd1+32];
	st.global.u64 	[%rd8+24], %rd13;
	.loc	4	73	0
	mov.u32 	%r13, 24;
	mov.u32 	%r2, 24;
	.loc	4	74	0
	mov.u32 	%r14, 30806;
	st.global.v2.u16 	[%rd8+32], {%r13,%r14};
	.loc	4	75	0
	st.global.u64 	[%rd1+24], %rd11;
	.loc	4	76	0
	st.global.u64 	[%rd1+32], %rd11;
	bra.uni 	$Lt_2_7426;
$Lt_2_7682:
	cvt.u32.u16 	%r2, %rh4;
$Lt_2_7426:
	.loc	4	79	0
	st.global.u16 	[%rd3+16], %r2;
	add.s32 	%r15, %r2, 16;
	mov.u32 	%r16, 1024;
	setp.le.s32 	%p5, %r15, %r16;
	@%p5 bra 	$Lt_2_7938;
	ld.global.u64 	%rd14, [%rd3+8];
	.loc	4	64	0
	ld.param.u64 	%rd1, [__cudaparm_AllocGet_t];
	.loc	4	79	0
	setp.ne.u64 	%p6, %rd2, %rd1;
	@%p6 bra 	$Lt_2_8706;
	.loc	4	64	0
	ld.param.u64 	%rd1, [__cudaparm_AllocGet_t];
	.loc	4	84	0
	st.global.u64 	[%rd1+32], %rd14;
	bra.uni 	$Lt_2_8450;
$Lt_2_8706:
	.loc	4	86	0
	st.global.u64 	[%rd2+8], %rd14;
$Lt_2_8450:
	.loc	4	87	0
	mov.u64 	%rd15, 0;
	st.global.u64 	[%rd3+8], %rd15;
$Lt_2_7938:
	.loc	18	17	0
	exit;
$LDWend_AllocGet:
	} // AllocGet

*/